#include <stdio.h>
#include <hip/hip_runtime.h>
#include <unistd.h>
#include <sys/select.h>

#define NSTREAM 2
#define BDIM 512

void initialData(float *ip, int size)
{
    int i;

    for(i = 0; i < size; i++)
    {
        ip[i] = (float)(rand() & 0xFF) / 10.0f;
    }
}

void sumArraysOnHost(float *A, float *B, float *C, const int N)
{
    for (int idx = 0; idx < N; idx++)
        C[idx] = A[idx] + B[idx];
}
__global__ __launch_bounds__(1024) 
void sumArrays(float *A, float *B, float *C, const int N)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    for (int i = 0; i < 1; ++i) {
    if (idx < N)
    {
        for(int j=0; j<60; j++)
        {
                C[idx] = A[idx] + B[idx] + i;
        }
    }
    }
}

void checkResult(float *hostRef, float *gpuRef, const int N)
{
    double epsilon = 1.0E-8;
    bool match = 1;

    for (int i = 0; i < N; i++)
    {
        if (abs(hostRef[i] - gpuRef[i]) > epsilon)
        {
            match = 0;
            printf("Arrays do not match!\n");
            printf("host %5.2f gpu %5.2f at %d\n", hostRef[i], gpuRef[i], i);
            break;
        }
    }

    if (match) printf("Arrays match.\n\n");
}

static void sleep_ms(unsigned int secs)
{
    struct timeval tval;
    tval.tv_sec = secs / 1000;
    tval.tv_usec = (secs * 1000) % 1000000;
    select(0, NULL, NULL, NULL, &tval);
}

int run_kernel()
{
     printf("> Starting...\n");

    int dev = 0;
    hipSetDevice(dev);
    hipDeviceProp_t deviceProp;
    hipGetDeviceProperties(&deviceProp, dev);
    printf("> Using Device %d: %s\n", dev, deviceProp.name);

    // set up data size of vectors
    //int nElem = 1 << 2;
    int nElem = 1 << 20;
    int iElem = nElem / NSTREAM;
    printf("> vector size = %d\n", nElem);
    size_t nBytes = nElem * sizeof(float);
    printf("> size nBytes = %ld\n", nBytes);

    hipStream_t stream0;
    hipStream_t stream1;
    hipStreamCreate(&stream0);
    hipStreamCreate(&stream1);
    // malloc pinned host memory for async memcpy
    float *h_A, *h_B, *h_C;
    hipHostMalloc((void**)&h_A, nBytes, hipHostMallocDefault);
    hipHostMalloc((void**)&h_B, nBytes, hipHostMallocDefault);
    hipHostMalloc((void**)&h_C, nBytes, hipHostMallocDefault);

    // initialize data at host side
    initialData(h_A, nElem);
    initialData(h_B, nElem);
    memset(h_C, 0, nBytes);

    // add vector at host side for result checks
    //sumArraysOnHost(h_A, h_B, hostRef, nElem);

    // malloc device global memory
    float *d_A, *d_B, *d_C;
    hipMalloc((float**)&d_A, nBytes);
    hipMalloc((float**)&d_B, nBytes);
    hipMalloc((float**)&d_C, nBytes);

    // invoke kernel at host side
    dim3 block (BDIM);
    dim3 grid  ((nElem + block.x - 1) / block.x);
    printf("> grid (%d,%d) block (%d,%d)\n", grid.x, grid.y, block.x, block.y);

    //warmup
    hipMemcpy(d_A, h_A, nBytes, hipMemcpyHostToDevice);
    hipMemcpy(d_B, h_B, nBytes, hipMemcpyHostToDevice);
    hipLaunchKernelGGL(sumArrays, dim3(grid), dim3(block), 0, stream0, d_A, d_B, d_C, nElem);
    hipMemcpy(h_C, d_C, nBytes, hipMemcpyDeviceToHost);

    hipFree(d_A);
    hipFree(d_B);
    hipFree(d_C);

    // free host memory
    hipHostFree(h_A);
    hipHostFree(h_B);
    hipHostFree(h_C);

    // destroy streams
    hipStreamDestroy(stream0);
    hipStreamDestroy(stream1);

    hipDeviceReset();
    printf("###END\n");
    return 0;
}
